#include <stdio.h>
#include <string.h>
#include "common.h"
#include "schema.h"
#include "gpulib.h"

__global__ static void merge(char * content, long * colOffset, int *outNum, int *attrSize, long *tupleNum, int *tupleSize, char *result){
	int startIndex = blockIdx.x*blockDim.x + threadIdx.x;
        long tNum = *tupleNum;
	int tSize = *tupleSize;
	int colNum = *outNum;
	
        int stride = blockDim.x * gridDim.x;
	long i;

	for(i=startIndex;i<tNum;i+=stride){
		int offset = 0;
		for(int j=0;j<colNum;j++){
			int index = j;
			int aSize = attrSize[index];
			memcpy(result+i*tSize + offset, content + colOffset[index]+ i*aSize, aSize);
			offset += aSize;
		}
	}
}

void * mergeCol(struct mergeNode * mn){
	struct tableNode *tn = mn->table;
	char * res, * gpuResult, *gpuContent;
	long size = tn->tupleNum * tn->tupleSize;
	int * gpuAttrSize, *gpuColNum, *gpuTupleSize;
	long * gpuColOffset, *colOffset, * gpuTupleNum;

	res = (char *) malloc(size);
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuResult, size));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuContent, size));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuColOffset,sizeof(long) * tn->totalAttr));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuColNum, sizeof(int)));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuAttrSize,sizeof(int) * tn->totalAttr));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuTupleNum, sizeof(long)));
	CUDA_SAFE_CALL_NO_SYNC(cudaMalloc((void **)&gpuTupleSize, sizeof(int)));

	long offset = 0;
	colOffset = (long *) malloc(sizeof(long) *tn->totalAttr);
	for(int i=0;i<tn->totalAttr;i++){
		colOffset[i] = offset;
		CUDA_SAFE_CALL_NO_SYNC(cudaMemcpy(gpuContent+offset, tn->content[i], tn->tupleNum *tn->attrSize[i], cudaMemcpyHostToDevice));
		offset += tn->tupleNum *tn->attrSize[i];
	}

	CUDA_SAFE_CALL(cudaMemcpy(gpuColOffset,colOffset, sizeof(long) * tn->totalAttr, cudaMemcpyHostToDevice));
	CUDA_SAFE_CALL(cudaMemcpy(gpuColNum, &(tn->totalAttr),sizeof(int), cudaMemcpyHostToDevice));
	CUDA_SAFE_CALL(cudaMemcpy(gpuAttrSize, tn->attrSize, sizeof(int) * tn->totalAttr, cudaMemcpyHostToDevice));
	CUDA_SAFE_CALL(cudaMemcpy(gpuTupleNum, &(tn->tupleNum),sizeof(long),cudaMemcpyHostToDevice));
	CUDA_SAFE_CALL(cudaMemcpy(gpuTupleSize, &(tn->tupleSize),sizeof(long),cudaMemcpyHostToDevice));

	dim3 grid(512);
	dim3 block(128);

	merge<<<grid,block>>> (gpuContent, gpuColOffset, gpuColNum, gpuAttrSize, gpuTupleNum, gpuTupleSize, gpuResult);

	CUDA_SAFE_CALL(cudaMemcpy(res, gpuResult, size, cudaMemcpyDeviceToHost));

	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuContent));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuColOffset));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuColNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuAttrSize));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuTupleNum));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuTupleSize));
	CUDA_SAFE_CALL_NO_SYNC(cudaFree(gpuResult));

	return res;
}
